





<!DOCTYPE html>
<html class="writer-html5" lang="zh-CN" >
<head>
  <meta charset="utf-8">
  
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <title>使用张量表达式来处理运算符 &mdash; tvm 0.8.dev1982 文档</title>
  

  
  <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/css/bootstrap.min.css" integrity="sha384-Gn5384xqQ1aoWXA+058RXPxPg6fy4IWvTNh0E263XmFcJlSAwiGgFAW/dAiS6JXm" crossorigin="anonymous">
  <link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../_static/gallery.css" type="text/css" />
  <link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../_static/css/tlcpack_theme.css" type="text/css" />

  
  
    <link rel="shortcut icon" href="../_static/tvm-logo-square.png"/>
  

  
  
  
  
    
      <script type="text/javascript" id="documentation_options" data-url_root="../" src="../_static/documentation_options.js"></script>
        <script data-url_root="../" id="documentation_options" src="../_static/documentation_options.js"></script>
        <script src="../_static/jquery.js"></script>
        <script src="../_static/underscore.js"></script>
        <script src="../_static/doctools.js"></script>
        <script src="../_static/translations.js"></script>
    
    <script type="text/javascript" src="../_static/js/theme.js"></script>

    
    <script type="text/javascript" src="../_static/js/tlcpack_theme.js"></script>
    <link rel="index" title="索引" href="../genindex.html" />
    <link rel="search" title="搜索" href="../search.html" />
    <link rel="next" title="Optimizing Operators with Schedule Templates and AutoTVM" href="autotvm_matmul_x86.html" />
    <link rel="prev" title="Compiling and Optimizing a Model with the Python Interface (AutoTVM)" href="autotvm_relay_x86.html" /> 
</head>

<body class="wy-body-for-nav">

   
  <div class="wy-grid-for-nav">
    
    
<header class="header">
    <div class="innercontainer">
      <div class="headerInner d-flex justify-content-between align-items-center">
          <div class="headerLogo">
               <a href="https://tvm.apache.org/"><img src=https://tvm.apache.org/assets/images/logo.svg alt="logo"></a>
          </div>

          <div id="headMenu" class="headerNav">
            <button type="button" id="closeHeadMenu" class="navCloseBtn"><img src="../_static/img/close-icon.svg" alt="Close"></button>
             <ul class="nav">
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/community>Community</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/download>Download</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/vta>VTA</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/blog>Blog</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/docs>Docs</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmconf.org>Conference</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://github.com/apache/tvm/>Github</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmchinese.github.io/declaration_zh_CN.html>About-Translators</a>
                </li>
             </ul>
               <div class="responsivetlcdropdown">
                 <button type="button" class="btn-link">
                   ASF
                 </button>
                 <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                 </ul>
               </div>
          </div>
            <div class="responsiveMenuIcon">
              <button type="button" id="menuBtn" class="btn-menu"><img src="../_static/img/menu-icon.svg" alt="Menu Icon"></button>
            </div>

            <div class="tlcDropdown">
              <div class="dropdown">
                <button type="button" class="btn-link dropdown-toggle" data-toggle="dropdown" aria-haspopup="true" aria-expanded="false">
                  ASF
                </button>
                <div class="dropdown-menu dropdown-menu-right">
                  <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                  </ul>
                </div>
              </div>
          </div>
       </div>
    </div>
 </header>
 
    <nav data-toggle="wy-nav-shift" class="wy-nav-side fixed">
      <div class="wy-side-scroll">
        <div class="wy-side-nav-search" >
          

          
            <a href="../index.html">
          

          
            
            <img src="../_static/tvm-logo-small.png" class="logo" alt="Logo"/>
          
          </a>

          
            
            
                <div class="version">
                  0.8.dev1982
                </div>
            
          

          
<div role="search">
  <form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
    <input type="hidden" name="check_keywords" value="yes" />
    <input type="hidden" name="area" value="default" />
  </form>
</div>

          
        </div>

        
        <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
          
            
            
              
            
            
              <p class="caption" role="heading"><span class="caption-text">如何开始</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../install/index.html">安装 TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../contribute/index.html">贡献者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">用户引导</span></p>
<ul class="current">
<li class="toctree-l1 current"><a class="reference internal" href="index.html">User Tutorial</a><ul class="current">
<li class="toctree-l2"><a class="reference internal" href="introduction.html">介绍</a></li>
<li class="toctree-l2"><a class="reference internal" href="introduction.html#an-overview-of-tvm-and-model-optimization">TVM和模型优化的概述</a></li>
<li class="toctree-l2"><a class="reference internal" href="install.html">安装 TVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="tvmc_command_line_driver.html">使用TVMC编译和优化一个模型</a></li>
<li class="toctree-l2"><a class="reference internal" href="autotvm_relay_x86.html">Compiling and Optimizing a Model with the Python Interface (AutoTVM)</a></li>
<li class="toctree-l2 current"><a class="current reference internal" href="#">使用张量表达式来处理运算符</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#example-1-writing-and-scheduling-vector-addition-in-te-for-cpu">例 1：为CPU用TE的形式写并调度一个向量加法。</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#describing-the-vector-computation">描述向量计算</a></li>
<li class="toctree-l4"><a class="reference internal" href="#create-a-default-schedule-for-the-computation">为计算创建一个默认的调度</a></li>
<li class="toctree-l4"><a class="reference internal" href="#compile-and-evaluate-the-default-schedule">编译并评估默认调度</a></li>
<li class="toctree-l4"><a class="reference internal" href="#updating-the-schedule-to-use-paralleism">更新调度来使用并行</a></li>
<li class="toctree-l4"><a class="reference internal" href="#updating-the-schedule-to-use-vectorization">更新调度以使用向量化</a></li>
<li class="toctree-l4"><a class="reference internal" href="#comparing-the-diferent-schedules">比较不同的调度</a></li>
<li class="toctree-l4"><a class="reference internal" href="#targeting-vector-addition-for-gpus-optional">TVM可以适配多种架构。在下一个示例中，我们将针对GPU编译向量加法。</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="#saving-and-loading-compiled-modules">保存和加载编译好的模块</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#load-compiled-module">加载编译好的模块</a></li>
<li class="toctree-l4"><a class="reference internal" href="#pack-everything-into-one-library">将所有内容打包到一个库中</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="#generate-opencl-code">产生OpenCL代码</a></li>
<li class="toctree-l3"><a class="reference internal" href="#example-2-manually-optimizing-matrix-multiplication-with-te">示例 2：使用 TE 手动优化矩阵乘法</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#preparation-and-performance-baseline">准备和基线（Baseline）</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-1-blocking">优化1：分块</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-2-vectorization">优化2：向量化</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-3-loop-permutation">优化3：循环重排</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-4-array-packing">优化4：数组压缩</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-5-optimizing-block-writing-through-caching">优化 5：通过缓存优化块写入</a></li>
<li class="toctree-l4"><a class="reference internal" href="#optimization-6-parallelization">优化 6：并行化</a></li>
<li class="toctree-l4"><a class="reference internal" href="#summary-of-matrix-multiplication-example">矩阵乘法示例总结</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="#final-notes-and-summary">最终说明总结</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="autotvm_matmul_x86.html">Optimizing Operators with Schedule Templates and AutoTVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="auto_scheduler_matmul_x86.html">Optimizing Operators with Auto-scheduling</a></li>
<li class="toctree-l2"><a class="reference internal" href="cross_compilation_and_rpc.html">Cross Compilation and RPC</a></li>
<li class="toctree-l2"><a class="reference internal" href="relay_quick_start.html">编译深度学习模型的快速开始教程</a></li>
<li class="toctree-l2"><a class="reference internal" href="intro_topi.html">Introduction to TOPI</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../how_to/index.html">How To Guides</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">开发者引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../dev/tutorial/index.html">Developer Tutorial</a></li>
<li class="toctree-l1"><a class="reference internal" href="../dev/how_to/how_to.html">开发者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">架构指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../arch/index.html">Design and Architecture</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">主题引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../topic/microtvm/index.html">microTVM：裸机使用TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../topic/vta/index.html">VTA: Versatile Tensor Accelerator</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">参考指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../reference/langref/index.html">语言参考</a></li>
<li class="toctree-l1"><a class="reference internal" href="../reference/api/python/index.html">Python API</a></li>
<li class="toctree-l1"><a class="reference internal" href="../reference/api/links.html">Other APIs</a></li>
<li class="toctree-l1"><a class="reference internal" href="../reference/publications.html">Publications</a></li>
<li class="toctree-l1"><a class="reference internal" href="../genindex.html">索引</a></li>
</ul>

            
          
        </div>
        
      </div>
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
      
      <nav class="wy-nav-top" aria-label="top navigation" data-toggle="wy-nav-top">
        
            <div class="togglemenu">

            </div>
            <div class="nav-content">
              <!-- tvm -->
              Table of content
            </div>
        
      </nav>


      <div class="wy-nav-content">
        
        <div class="rst-content">
        

          




















<div role="navigation" aria-label="breadcrumbs navigation">

  <ul class="wy-breadcrumbs">
    
      <li><a href="../index.html">Docs</a> <span class="br-arrow">></span></li>
        
          <li><a href="index.html">User Tutorial</a> <span class="br-arrow">></span></li>
        
      <li>使用张量表达式来处理运算符</li>
    
    
      <li class="wy-breadcrumbs-aside">
        
            
            <a href="../_sources/tutorial/tensor_expr_get_started.rst.txt" rel="nofollow"> <img src="../_static//img/source.svg" alt="viewsource"/></a>
          
        
      </li>
    
  </ul>

  
  <hr/>
</div>
          <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
           <div itemprop="articleBody">
            
  <div class="sphx-glr-download-link-note admonition note">
<p class="admonition-title">注解</p>
<p>点击 <a class="reference internal" href="#sphx-glr-download-tutorial-tensor-expr-get-started-py"><span class="std std-ref">这里</span></a> 下载完整的样例代码</p>
</div>
<div class="sphx-glr-example-title section" id="working-with-operators-using-tensor-expression">
<span id="tutorial-tensor-expr-get-started"></span><span id="sphx-glr-tutorial-tensor-expr-get-started-py"></span><h1>使用张量表达式来处理运算符<a class="headerlink" href="#working-with-operators-using-tensor-expression" title="永久链接至标题">¶</a></h1>
<p><strong>作者</strong>: <a class="reference external" href="https://tqchen.github.io">Tianqi Chen</a></p>
<p>本次教程，我们会关注于TVM怎样使用张量表达式（TE）来定义张量计算和优化循环。 TE 用一种纯函数语言来描述张量计算（这意味着每一个表达式都没有副作用）。当在TVM整体环境下来看，Relay 用一系列操作符来描述一个计算过程。这一系列操作符中的每一个操作符可以都被表示为一个TE表达式，并且这每一个TE表达式都采用张量输入，然后生成一个张量输出。</p>
<p>这是TVM中的张量表达式语言的介绍性教程。TVM用一种领域特定张量表达式来进行高效的内核构建。我们会用两个使用张量表达式语言的例子来演示基本的工作流程。第一个例子我们会用向量加法来介绍TE和调度。第二个例子用TE进行了一个矩阵乘法的逐步优化，来拓展了这两个概念。这个矩阵乘法的示例会被当做未来关于TVM更多高级功能的教程的比较基础。</p>
<div class="section" id="example-1-writing-and-scheduling-vector-addition-in-te-for-cpu">
<h2>例 1：为CPU用TE的形式写并调度一个向量加法。<a class="headerlink" href="#example-1-writing-and-scheduling-vector-addition-in-te-for-cpu" title="永久链接至标题">¶</a></h2>
<p>让我们来看一个python示例，我们会实现用一个TE来表达向量加法，然后针对CPU做一个调度。我们首先初始化TVM的环境。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">tvm</span>
<span class="kn">import</span> <span class="nn">tvm.testing</span>
<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">te</span>
<span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span>

<span class="c1"># You will get better performance if you can identify the CPU you are targeting</span>
<span class="c1"># and specify it. If you&#39;re using llvm, you can get this information from the</span>
<span class="c1"># command ``llc --version`` to get the CPU type, and you can check</span>
<span class="c1"># ``/proc/cpuinfo`` for additional extensions that your processor might</span>
<span class="c1"># support. For example, you can use &quot;llvm -mcpu=skylake-avx512&quot; for CPUs with</span>
<span class="c1"># AVX-512 instructions.</span>

<span class="n">tgt</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span><span class="p">(</span><span class="n">target</span><span class="o">=</span><span class="s2">&quot;llvm&quot;</span><span class="p">,</span> <span class="n">host</span><span class="o">=</span><span class="s2">&quot;llvm&quot;</span><span class="p">)</span>
</pre></div>
</div>
<div class="section" id="describing-the-vector-computation">
<h3>描述向量计算<a class="headerlink" href="#describing-the-vector-computation" title="永久链接至标题">¶</a></h3>
<p>我们描述一个向量加法计算。TVM采用张量语义，将每一个中间结果都表示成一个多维数组。用户需要描述生成张量的计算规则。我们首先定义一个符号变量 <code class="docutils literal notranslate"><span class="pre">n</span></code> 来表示shape。然后我们来定义两个占位张量 <code class="docutils literal notranslate"><span class="pre">A</span></code> 和 <code class="docutils literal notranslate"><span class="pre">B</span></code>，他们具有相同的shape <code class="docutils literal notranslate"><span class="pre">(n,)</span></code>。然后我们用一个 <code class="docutils literal notranslate"><span class="pre">compute</span></code> 操作来描述结果张量``C``。 <code class="docutils literal notranslate"><span class="pre">compute</span></code> 定义了一个计算，输出张量符合指定的形状， 并在由lambda函数定义的张量中的每个位置执行计算。请注意，虽然 <code class="docutils literal notranslate"><span class="pre">n</span></code> 是一个变量，但它定义了 <code class="docutils literal notranslate"><span class="pre">A</span></code>, <code class="docutils literal notranslate"><span class="pre">B</span></code> 和``C`` 张量之间统一的形状。请记住，在此阶段不会有任何实际的计算发生，因为我们只是声明了这个计算该怎么被执行。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">n</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">var</span><span class="p">(</span><span class="s2">&quot;n&quot;</span><span class="p">)</span>
<span class="n">A</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;A&quot;</span><span class="p">)</span>
<span class="n">B</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;B&quot;</span><span class="p">)</span>
<span class="n">C</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">shape</span><span class="p">,</span> <span class="k">lambda</span> <span class="n">i</span><span class="p">:</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">+</span> <span class="n">B</span><span class="p">[</span><span class="n">i</span><span class="p">],</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;C&quot;</span><span class="p">)</span>
</pre></div>
</div>
<div class="admonition note">
<p class="admonition-title">注解</p>
<p>Lambda函数</p>
<p>The second argument to the <code class="docutils literal notranslate"><span class="pre">te.compute</span></code> method is the function that
performs the computation. In this example, we’re using an anonymous function,
also known as a <code class="docutils literal notranslate"><span class="pre">lambda</span></code> function, to define the computation, in this case
addition on the <code class="docutils literal notranslate"><span class="pre">i``th</span> <span class="pre">element</span> <span class="pre">of</span> <span class="pre">``A</span></code> and <code class="docutils literal notranslate"><span class="pre">B</span></code>.</p>
</div>
</div>
<div class="section" id="create-a-default-schedule-for-the-computation">
<h3>为计算创建一个默认的调度<a class="headerlink" href="#create-a-default-schedule-for-the-computation" title="永久链接至标题">¶</a></h3>
<p>虽然上文定义了计算的规则，我们可以用不同的方式来计算``C``以适应不同的设备。对于一个多轴张量来说，你可以选择先迭代哪个轴，或者可以跨不同的线程来拆分计算。TVM需要用户提供一个调度，即对计算该怎样执行的描述。TE中的调度操作可以改变循环的顺序，跨不同线程拆分计算，数据分块计算等等。调度背后的一个重要概念是它们仅仅描述计算是如何被执行的，所以对一个相同的TE来说，不同的调度也会生成相同的结果。</p>
<p>TVM允许你创建一个简单的调度，按行优先迭代顺序来计算 <code class="docutils literal notranslate"><span class="pre">C</span></code>。</p>
<div class="highlight-c notranslate"><div class="highlight"><pre><span></span><span class="k">for</span> <span class="p">(</span><span class="kt">int</span> <span class="n">i</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">i</span> <span class="o">&lt;</span> <span class="n">n</span><span class="p">;</span> <span class="o">++</span><span class="n">i</span><span class="p">)</span> <span class="p">{</span>
  <span class="n">C</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">=</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">+</span> <span class="n">B</span><span class="p">[</span><span class="n">i</span><span class="p">];</span>
<span class="p">}</span>
</pre></div>
</div>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>
</pre></div>
</div>
</div>
<div class="section" id="compile-and-evaluate-the-default-schedule">
<h3>编译并评估默认调度<a class="headerlink" href="#compile-and-evaluate-the-default-schedule" title="永久链接至标题">¶</a></h3>
<p>有了TE表达式和调度，我们就能为目标语言和架构（在本例中为LLVM和CPU）生成可运行的代码。我们向TVM提供调度，使用该调度的TE表达式列表，目标和主机，以及我们生成的函数名称。输出的结果是可以直接从Python调用的一个类型被擦除的函数。</p>
<p>接下来这行，我们用tvm.build来创建一个函数。build函数接收调度，所需的函数签名（包括输入和输出）还有我们想要编译成的目标语言。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">fadd</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">tgt</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;myadd&quot;</span><span class="p">)</span>
</pre></div>
</div>
<p>让我们来运行这个函数，然后将输出与numpy中相同的计算结果进行比较。这个编译后的TVM函数被暴露为一个简洁的C API，可以用任何语言来调用。我们首先创建一个TVM可以将调度编译上去的设备，在本例中是LLVM CPU。然后我们可以在设备中初始化张量，再执行我们自定义的计算。为了验证计算是否正确，我们可以将C张量的输出结果跟numpy执行的相同计算的结果进行比较。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">device</span><span class="p">(</span><span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="p">,</span> <span class="mi">0</span><span class="p">)</span>

<span class="n">n</span> <span class="o">=</span> <span class="mi">1024</span>
<span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">B</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">C</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">fadd</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>
</pre></div>
</div>
<p>为了比较这种形式的计算和numpy计算的速度，创建一个辅助函数来运行TVM生成的代码生成一个profile文件。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">timeit</span>

<span class="n">np_repeat</span> <span class="o">=</span> <span class="mi">100</span>
<span class="n">np_running_time</span> <span class="o">=</span> <span class="n">timeit</span><span class="o">.</span><span class="n">timeit</span><span class="p">(</span>
    <span class="n">setup</span><span class="o">=</span><span class="s2">&quot;import numpy</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;n = 32768</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s1">&#39;dtype = &quot;float32&quot;</span><span class="se">\n</span><span class="s1">&#39;</span>
    <span class="s2">&quot;a = numpy.random.rand(n, 1).astype(dtype)</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;b = numpy.random.rand(n, 1).astype(dtype)</span><span class="se">\n</span><span class="s2">&quot;</span><span class="p">,</span>
    <span class="n">stmt</span><span class="o">=</span><span class="s2">&quot;answer = a + b&quot;</span><span class="p">,</span>
    <span class="n">number</span><span class="o">=</span><span class="n">np_repeat</span><span class="p">,</span>
<span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Numpy running time: </span><span class="si">%f</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="n">np_running_time</span> <span class="o">/</span> <span class="n">np_repeat</span><span class="p">))</span>


<span class="k">def</span> <span class="nf">evaluate_addition</span><span class="p">(</span><span class="n">func</span><span class="p">,</span> <span class="n">target</span><span class="p">,</span> <span class="n">optimization</span><span class="p">,</span> <span class="n">log</span><span class="p">):</span>
    <span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">device</span><span class="p">(</span><span class="n">target</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="p">,</span> <span class="mi">0</span><span class="p">)</span>
    <span class="n">n</span> <span class="o">=</span> <span class="mi">32768</span>
    <span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">B</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">C</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>

    <span class="n">evaluator</span> <span class="o">=</span> <span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span><span class="p">(</span><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">number</span><span class="o">=</span><span class="mi">10</span><span class="p">)</span>
    <span class="n">mean_time</span> <span class="o">=</span> <span class="n">evaluator</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span><span class="o">.</span><span class="n">mean</span>
    <span class="nb">print</span><span class="p">(</span><span class="s2">&quot;</span><span class="si">%s</span><span class="s2">: </span><span class="si">%f</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="n">optimization</span><span class="p">,</span> <span class="n">mean_time</span><span class="p">))</span>

    <span class="n">log</span><span class="o">.</span><span class="n">append</span><span class="p">((</span><span class="n">optimization</span><span class="p">,</span> <span class="n">mean_time</span><span class="p">))</span>


<span class="n">log</span> <span class="o">=</span> <span class="p">[(</span><span class="s2">&quot;numpy&quot;</span><span class="p">,</span> <span class="n">np_running_time</span> <span class="o">/</span> <span class="n">np_repeat</span><span class="p">)]</span>
<span class="n">evaluate_addition</span><span class="p">(</span><span class="n">fadd</span><span class="p">,</span> <span class="n">tgt</span><span class="p">,</span> <span class="s2">&quot;naive&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Numpy running time: 0.000008
naive: 0.000008
</pre></div>
</div>
</div>
<div class="section" id="updating-the-schedule-to-use-paralleism">
<h3>更新调度来使用并行<a class="headerlink" href="#updating-the-schedule-to-use-paralleism" title="永久链接至标题">¶</a></h3>
<p>现在我们已经说明了TE的基础知识，让我们更深度地了解调度的作用，以及如何使用它们来优化不同加架构的张量表达式。调度是应用于将一个表达式通过多种不同的方式进行变换的一系列步骤。当调度应用于TE中的表达式时，输入和输出保持不变，但当编译时表达式的实现可能会发生变化。在默认的调度中，这种张量加运算是串行的，但很容易跨所有处理器线程并行话。我们可以将并行调度操作应用于我们的计算。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">])</span>
</pre></div>
</div>
<p><a href="#id1"><span class="problematic" id="id2">``</span></a>tvm.lower``命令将生成TE的中间表示（IR），以及对应的调度。通过lower我们应用不同的调度操作的表达式，我们可以看到调度对计算顺序的影响。我们使用``simple_mode=True``标志来返回一个可读的C风格语句。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [n: int32], [stride: int32], type=&quot;auto&quot;),
             A: Buffer(A_2: Pointer(float32), float32, [n], [stride_1: int32], type=&quot;auto&quot;),
             B: Buffer(B_2: Pointer(float32), float32, [n], [stride_2: int32], type=&quot;auto&quot;)}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (i: int32, 0, n) &quot;parallel&quot; {
    C_2[(i*stride)] = ((float32*)A_2[(i*stride_1)] + (float32*)B_2[(i*stride_2)])
  }
}
</pre></div>
</div>
<p>TVM现在可以在独立的线程中运行这些块。让我们编译和运行这个新的应用了并行操作的调度：</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">fadd_parallel</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">tgt</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;myadd_parallel&quot;</span><span class="p">)</span>
<span class="n">fadd_parallel</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>

<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>

<span class="n">evaluate_addition</span><span class="p">(</span><span class="n">fadd_parallel</span><span class="p">,</span> <span class="n">tgt</span><span class="p">,</span> <span class="s2">&quot;parallel&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>parallel: 0.000008
</pre></div>
</div>
</div>
<div class="section" id="updating-the-schedule-to-use-vectorization">
<h3>更新调度以使用向量化<a class="headerlink" href="#updating-the-schedule-to-use-vectorization" title="永久链接至标题">¶</a></h3>
<p>现代CPU还能够对浮点值执行SIMD操作，我们可以通过将另外一个调度应用到我们的表达式计算中以利用这一点。实现这个需要多个步骤：首先，我们必须使用split调度原语将调度拆分为内循环和外循环。内循环可以使用矢量化调度源语来使用SIMD指令进行矢量化，外循环可以使用并行调度源语并行化。选择拆分因子为你CPU上的线程数。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># Recreate the schedule, since we modified it with the parallel operation in</span>
<span class="c1"># the previous example</span>
<span class="n">n</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">var</span><span class="p">(</span><span class="s2">&quot;n&quot;</span><span class="p">)</span>
<span class="n">A</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;A&quot;</span><span class="p">)</span>
<span class="n">B</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;B&quot;</span><span class="p">)</span>
<span class="n">C</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">shape</span><span class="p">,</span> <span class="k">lambda</span> <span class="n">i</span><span class="p">:</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">+</span> <span class="n">B</span><span class="p">[</span><span class="n">i</span><span class="p">],</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;C&quot;</span><span class="p">)</span>

<span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>

<span class="c1"># This factor should be chosen to match the number of threads appropriate for</span>
<span class="c1"># your CPU. This will vary depending on architecture, but a good rule is</span>
<span class="c1"># setting this factor to equal the number of available CPU cores.</span>
<span class="n">factor</span> <span class="o">=</span> <span class="mi">4</span>

<span class="n">outer</span><span class="p">,</span> <span class="n">inner</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">factor</span><span class="o">=</span><span class="n">factor</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">outer</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">inner</span><span class="p">)</span>

<span class="n">fadd_vector</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">tgt</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;myadd_parallel&quot;</span><span class="p">)</span>

<span class="n">evaluate_addition</span><span class="p">(</span><span class="n">fadd_vector</span><span class="p">,</span> <span class="n">tgt</span><span class="p">,</span> <span class="s2">&quot;vector&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>

<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>vector: 0.000035
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [n: int32], [stride: int32], type=&quot;auto&quot;),
             C: Buffer(C_2: Pointer(float32), float32, [n], [stride_1: int32], type=&quot;auto&quot;),
             B: Buffer(B_2: Pointer(float32), float32, [n], [stride_2: int32], type=&quot;auto&quot;)}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (i.outer: int32, 0, floordiv((n + 3), 4)) &quot;parallel&quot; {
    for (i.inner.s: int32, 0, 4) {
      if @tir.likely((((i.outer*4) + i.inner.s) &lt; n), dtype=bool) {
        C_2[(((i.outer*4) + i.inner.s)*stride_1)] = ((float32*)A_2[(((i.outer*4) + i.inner.s)*stride)] + (float32*)B_2[(((i.outer*4) + i.inner.s)*stride_2)])
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="comparing-the-diferent-schedules">
<h3>比较不同的调度<a class="headerlink" href="#comparing-the-diferent-schedules" title="永久链接至标题">¶</a></h3>
<p>我们现在可以比较不同的调度</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">baseline</span> <span class="o">=</span> <span class="n">log</span><span class="p">[</span><span class="mi">0</span><span class="p">][</span><span class="mi">1</span><span class="p">]</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="s2">&quot;Operator&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="s2">&quot;Timing&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="s2">&quot;Performance&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">)))</span>
<span class="k">for</span> <span class="n">result</span> <span class="ow">in</span> <span class="n">log</span><span class="p">:</span>
    <span class="nb">print</span><span class="p">(</span>
        <span class="s2">&quot;</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="s2">&quot;</span>
        <span class="o">%</span> <span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="nb">str</span><span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">1</span><span class="p">])</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="nb">str</span><span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">1</span><span class="p">]</span> <span class="o">/</span> <span class="n">baseline</span><span class="p">)</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">))</span>
    <span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Operator                  Timing             Performance
   numpy    8.262577466666698e-06                    1.0
   naive              8.0611e-06       0.975615663819249
parallel              7.5443e-06      0.9130685951733089
  vector             3.45197e-05       4.177836775420393
</pre></div>
</div>
<div class="admonition note">
<p class="admonition-title">注解</p>
<p>代码定制化</p>
<p>你可能已经注意到，<code class="docutils literal notranslate"><span class="pre">A</span></code>, <code class="docutils literal notranslate"><span class="pre">B</span></code> 和``C`` 的声明都采用相同的形状参数 <code class="docutils literal notranslate"><span class="pre">n</span></code>。TVM将利用这一点将单个形状参数传递给内核，正如你在打印的设备代码中可以找到的那样。这是定制化的一种形式。</p>
<p>在主机端，TVM会自动生成检查代码，检查参数中的约束。因此，如果将具有不同形状的数组传递给fadd，则会出错。</p>
<p>我们可以做更多的定制化。例如，我们可以在计算声明中写 <code class="code docutils literal notranslate"><span class="pre">n</span> <span class="pre">=</span> <span class="pre">tvm.runtime.convert(1024)</span></code> 而不是 <code class="code docutils literal notranslate"><span class="pre">n</span> <span class="pre">=</span> <span class="pre">te.var(&quot;n&quot;)</span></code> 。生成的函数将只接收长度为1024的向量。</p>
</div>
<p>我们已经定义，调度和编译了一个向量加法运算符，然后我们就可以在TVM运行时执行它。我们课可以将运算符保存为一个库，然后我们可以使用TVM运行时在稍后进行加载。</p>
</div>
<div class="section" id="targeting-vector-addition-for-gpus-optional">
<h3>TVM可以适配多种架构。在下一个示例中，我们将针对GPU编译向量加法。<a class="headerlink" href="#targeting-vector-addition-for-gpus-optional" title="永久链接至标题">¶</a></h3>
<p>TVM可以适配多种架构。在下一个例子中，我们将基于GPU编译矢量加法。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># If you want to run this code, change ``run_cuda = True``</span>
<span class="c1"># Note that by default this example is not run in the docs CI.</span>

<span class="n">run_cuda</span> <span class="o">=</span> <span class="kc">False</span>
<span class="k">if</span> <span class="n">run_cuda</span><span class="p">:</span>
    <span class="c1"># Change this target to the correct backend for you gpu. For example: cuda (NVIDIA GPUs),</span>
    <span class="c1"># rocm (Radeon GPUS), OpenCL (opencl).</span>
    <span class="n">tgt_gpu</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span><span class="p">(</span><span class="n">target</span><span class="o">=</span><span class="s2">&quot;cuda&quot;</span><span class="p">,</span> <span class="n">host</span><span class="o">=</span><span class="s2">&quot;llvm&quot;</span><span class="p">)</span>

    <span class="c1"># Recreate the schedule</span>
    <span class="n">n</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">var</span><span class="p">(</span><span class="s2">&quot;n&quot;</span><span class="p">)</span>
    <span class="n">A</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;A&quot;</span><span class="p">)</span>
    <span class="n">B</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">n</span><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;B&quot;</span><span class="p">)</span>
    <span class="n">C</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">shape</span><span class="p">,</span> <span class="k">lambda</span> <span class="n">i</span><span class="p">:</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">+</span> <span class="n">B</span><span class="p">[</span><span class="n">i</span><span class="p">],</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;C&quot;</span><span class="p">)</span>
    <span class="nb">print</span><span class="p">(</span><span class="nb">type</span><span class="p">(</span><span class="n">C</span><span class="p">))</span>

    <span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>

    <span class="n">bx</span><span class="p">,</span> <span class="n">tx</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">factor</span><span class="o">=</span><span class="mi">64</span><span class="p">)</span>

    <span class="c1">################################################################################</span>
    <span class="c1"># Finally we must bind the iteration axis bx and tx to threads in the GPU</span>
    <span class="c1"># compute grid. The naive schedule is not valid for GPUs, and these are</span>
    <span class="c1"># specific constructs that allow us to generate code that runs on a GPU.</span>

    <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">bx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.x&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.x&quot;</span><span class="p">))</span>

    <span class="c1">######################################################################</span>
    <span class="c1"># Compilation</span>
    <span class="c1"># -----------</span>
    <span class="c1"># After we have finished specifying the schedule, we can compile it</span>
    <span class="c1"># into a TVM function. By default TVM compiles into a type-erased</span>
    <span class="c1"># function that can be directly called from the python side.</span>
    <span class="c1">#</span>
    <span class="c1"># In the following line, we use tvm.build to create a function.</span>
    <span class="c1"># The build function takes the schedule, the desired signature of the</span>
    <span class="c1"># function (including the inputs and outputs) as well as target language</span>
    <span class="c1"># we want to compile to.</span>
    <span class="c1">#</span>
    <span class="c1"># The result of compilation fadd is a GPU device function (if GPU is</span>
    <span class="c1"># involved) as well as a host wrapper that calls into the GPU</span>
    <span class="c1"># function. fadd is the generated host wrapper function, it contains</span>
    <span class="c1"># a reference to the generated device function internally.</span>

    <span class="n">fadd</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">tgt_gpu</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;myadd&quot;</span><span class="p">)</span>

    <span class="c1">################################################################################</span>
    <span class="c1"># The compiled TVM function is exposes a concise C API that can be invoked from</span>
    <span class="c1"># any language.</span>
    <span class="c1">#</span>
    <span class="c1"># We provide a minimal array API in python to aid quick testing and prototyping.</span>
    <span class="c1"># The array API is based on the `DLPack &lt;https://github.com/dmlc/dlpack&gt;`_ standard.</span>
    <span class="c1">#</span>
    <span class="c1"># - We first create a GPU device.</span>
    <span class="c1"># - Then tvm.nd.array copies the data to the GPU.</span>
    <span class="c1"># - ``fadd`` runs the actual computation</span>
    <span class="c1"># - ``numpy()`` copies the GPU array back to the CPU (so we can verify correctness).</span>
    <span class="c1">#</span>
    <span class="c1"># Note that copying the data to and from the memory on the GPU is a required step.</span>

    <span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">device</span><span class="p">(</span><span class="n">tgt_gpu</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="p">,</span> <span class="mi">0</span><span class="p">)</span>

    <span class="n">n</span> <span class="o">=</span> <span class="mi">1024</span>
    <span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">B</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">C</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">fadd</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
    <span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>

    <span class="c1">################################################################################</span>
    <span class="c1"># Inspect the Generated GPU Code</span>
    <span class="c1"># ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~</span>
    <span class="c1"># You can inspect the generated code in TVM. The result of tvm.build is a TVM</span>
    <span class="c1"># Module. fadd is the host module that contains the host wrapper, it also</span>
    <span class="c1"># contains a device module for the CUDA (GPU) function.</span>
    <span class="c1">#</span>
    <span class="c1"># The following code fetches the device module and prints the content code.</span>

    <span class="k">if</span> <span class="p">(</span>
        <span class="n">tgt_gpu</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;cuda&quot;</span>
        <span class="ow">or</span> <span class="n">tgt_gpu</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;rocm&quot;</span>
        <span class="ow">or</span> <span class="n">tgt_gpu</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="o">.</span><span class="n">startswith</span><span class="p">(</span><span class="s2">&quot;opencl&quot;</span><span class="p">)</span>
    <span class="p">):</span>
        <span class="n">dev_module</span> <span class="o">=</span> <span class="n">fadd</span><span class="o">.</span><span class="n">imported_modules</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span>
        <span class="nb">print</span><span class="p">(</span><span class="s2">&quot;-----GPU code-----&quot;</span><span class="p">)</span>
        <span class="nb">print</span><span class="p">(</span><span class="n">dev_module</span><span class="o">.</span><span class="n">get_source</span><span class="p">())</span>
    <span class="k">else</span><span class="p">:</span>
        <span class="nb">print</span><span class="p">(</span><span class="n">fadd</span><span class="o">.</span><span class="n">get_source</span><span class="p">())</span>
</pre></div>
</div>
</div>
</div>
<div class="section" id="saving-and-loading-compiled-modules">
<h2>保存和加载编译好的模块<a class="headerlink" href="#saving-and-loading-compiled-modules" title="永久链接至标题">¶</a></h2>
<p>除了运行时编译之外，我们还可以将编译后的模块保存到一个文件中，稍后再加载它们。</p>
<p>以下代码首先执行以下步骤：</p>
<ul class="simple">
<li><p>它将编译后的host模块保存到一个目标文件中。</p></li>
<li><p>然后它将device模块保存到一个 ptx 文件中。</p></li>
<li><p>cc.create_shared 调用编译器（gcc）来创建共享库</p></li>
</ul>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">from</span> <span class="nn">tvm.contrib</span> <span class="k">import</span> <span class="n">cc</span>
<span class="kn">from</span> <span class="nn">tvm.contrib</span> <span class="k">import</span> <span class="n">utils</span>

<span class="n">temp</span> <span class="o">=</span> <span class="n">utils</span><span class="o">.</span><span class="n">tempdir</span><span class="p">()</span>
<span class="n">fadd</span><span class="o">.</span><span class="n">save</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.o&quot;</span><span class="p">))</span>
<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;cuda&quot;</span><span class="p">:</span>
    <span class="n">fadd</span><span class="o">.</span><span class="n">imported_modules</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">save</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.ptx&quot;</span><span class="p">))</span>
<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;rocm&quot;</span><span class="p">:</span>
    <span class="n">fadd</span><span class="o">.</span><span class="n">imported_modules</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">save</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.hsaco&quot;</span><span class="p">))</span>
<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="o">.</span><span class="n">startswith</span><span class="p">(</span><span class="s2">&quot;opencl&quot;</span><span class="p">):</span>
    <span class="n">fadd</span><span class="o">.</span><span class="n">imported_modules</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">save</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.cl&quot;</span><span class="p">))</span>
<span class="n">cc</span><span class="o">.</span><span class="n">create_shared</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.so&quot;</span><span class="p">),</span> <span class="p">[</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.o&quot;</span><span class="p">)])</span>
<span class="nb">print</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">listdir</span><span class="p">())</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>[&#39;myadd.o&#39;, &#39;myadd.so&#39;]
</pre></div>
</div>
<div class="admonition note">
<p class="admonition-title">注解</p>
<p>模块存储格式</p>
<p>CPU（主机）模块直接保存为共享库（.so）。 设备代码可以有多种自定义格式。 在我们的示例中，设备代码存储在 ptx 中，也是一个元数据 json 文件。 它们可以通过导入单独加载和链接。</p>
</div>
<div class="section" id="load-compiled-module">
<h3>加载编译好的模块<a class="headerlink" href="#load-compiled-module" title="永久链接至标题">¶</a></h3>
<p>我们可以从文件系统加载编译好的模块并运行代码。 以下代码分别加载主机和设备模块并将它们链接在一起。 我们可以验证新加载的函数是否有效。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">fadd1</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">load_module</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.so&quot;</span><span class="p">))</span>
<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;cuda&quot;</span><span class="p">:</span>
    <span class="n">fadd1_dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">load_module</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.ptx&quot;</span><span class="p">))</span>
    <span class="n">fadd1</span><span class="o">.</span><span class="n">import_module</span><span class="p">(</span><span class="n">fadd1_dev</span><span class="p">)</span>

<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">&quot;rocm&quot;</span><span class="p">:</span>
    <span class="n">fadd1_dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">load_module</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.hsaco&quot;</span><span class="p">))</span>
    <span class="n">fadd1</span><span class="o">.</span><span class="n">import_module</span><span class="p">(</span><span class="n">fadd1_dev</span><span class="p">)</span>

<span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="o">.</span><span class="n">startswith</span><span class="p">(</span><span class="s2">&quot;opencl&quot;</span><span class="p">):</span>
    <span class="n">fadd1_dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">load_module</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd.cl&quot;</span><span class="p">))</span>
    <span class="n">fadd1</span><span class="o">.</span><span class="n">import_module</span><span class="p">(</span><span class="n">fadd1_dev</span><span class="p">)</span>

<span class="n">fadd1</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>
</pre></div>
</div>
</div>
<div class="section" id="pack-everything-into-one-library">
<h3>将所有内容打包到一个库中<a class="headerlink" href="#pack-everything-into-one-library" title="永久链接至标题">¶</a></h3>
<p>在上面的例子中，我们分别存储设备和主机代码。 TVM 还支持将所有内容导出为一个共享库。 在底层，我们将设备模块打包成二进制 blob，并将它们与主机代码链接在一起。 目前我们支持 Metal、OpenCL 和 CUDA 模块的打包。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">fadd</span><span class="o">.</span><span class="n">export_library</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd_pack.so&quot;</span><span class="p">))</span>
<span class="n">fadd2</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">load_module</span><span class="p">(</span><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span><span class="p">(</span><span class="s2">&quot;myadd_pack.so&quot;</span><span class="p">))</span>
<span class="n">fadd2</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>
</pre></div>
</div>
<div class="admonition note">
<p class="admonition-title">注解</p>
<p>运行时API和线程安全</p>
<p>TVM的编译模块不依赖于TVM编译器。相反，它只依赖于最小的运行时库。TVM运行时库包装了设备驱动程序，并为编译的函数提供了线程安全和设备无关的调用。</p>
<p>这意味着你可以从任何线程，任何GPU上调用已编译的TVM函数，前提是你已经为该GPU编译了代码。</p>
</div>
</div>
</div>
<div class="section" id="generate-opencl-code">
<h2>产生OpenCL代码<a class="headerlink" href="#generate-opencl-code" title="永久链接至标题">¶</a></h2>
<p>TVM 向多个后端提供代码生成功能。 我们还可以生成在 CPU 后端运行的 OpenCL 代码或 LLVM 代码。</p>
<p>以下代码块生成 OpenCL 代码，在 OpenCL 设备上创建数组，并验证代码的正确性。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="k">if</span> <span class="n">tgt</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="o">.</span><span class="n">startswith</span><span class="p">(</span><span class="s2">&quot;opencl&quot;</span><span class="p">):</span>
    <span class="n">fadd_cl</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">tgt</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;myadd&quot;</span><span class="p">)</span>
    <span class="nb">print</span><span class="p">(</span><span class="s2">&quot;------opencl code------&quot;</span><span class="p">)</span>
    <span class="nb">print</span><span class="p">(</span><span class="n">fadd_cl</span><span class="o">.</span><span class="n">imported_modules</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">get_source</span><span class="p">())</span>
    <span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">cl</span><span class="p">(</span><span class="mi">0</span><span class="p">)</span>
    <span class="n">n</span> <span class="o">=</span> <span class="mi">1024</span>
    <span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="n">n</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">B</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">C</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">fadd_cl</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
    <span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>
</pre></div>
</div>
<div class="admonition note">
<p class="admonition-title">注解</p>
<p>TE调度源语</p>
<p>TVM包括许多不同的调度源语：</p>
<ul class="simple">
<li><p>split：将指定的轴按定义的因子拆分为两个轴。</p></li>
<li><p>tile：tile将按照定义的因子将计算拆分到两个轴上。</p></li>
<li><p>fuse：融合一个计算的两个连续轴。</p></li>
<li><p>reorder：可以将计算的轴重排为定义的顺序。</p></li>
<li><p>bind：可以将计算绑定到特定线程，这在 GPU 编程中很有用。</p></li>
<li><p>compute_at：默认情况下，TVM 将在函数的最外层或根计算张量。 compute_at 指定应该在另一个算子的第一个计算轴上计算一个张量。</p></li>
<li><p>compute_inline：当标记为内联时，计算将被扩展然后插入到被张量需要的地址中。</p></li>
<li><p>compute_root：将计算移动到函数的最外层或根。 这意味着在进入下一个阶段之前，计算的阶段将被完全计算。</p></li>
</ul>
<p>这些原语的完整描述可以在 [调度源语](<a class="reference external" href="https://tvm.apache.org/docs/tutorials/language/schedule_primitives.html">https://tvm.apache.org/docs/tutorials/language/schedule_primitives.html</a>) 文档页面中找到。</p>
</div>
</div>
<div class="section" id="example-2-manually-optimizing-matrix-multiplication-with-te">
<h2>示例 2：使用 TE 手动优化矩阵乘法<a class="headerlink" href="#example-2-manually-optimizing-matrix-multiplication-with-te" title="永久链接至标题">¶</a></h2>
<p>现在我们将考虑第二个更高级的示例，演示如何仅使用 18 行 Python 代码 TVM 将常见矩阵乘法运算速度提高 18 倍。</p>
<p><strong>矩阵乘法是计算密集型操作。要实现良好的CPU性能有两个重要的优化：</strong></p>
<ol class="arabic simple">
<li><p>提高内存访问的缓存命中率。复杂的数值计算和热点内存访问都可以通过高缓存命中率来加速。 这需要我们将原始内存访问模式转换为适合缓存策略的模式。</p></li>
<li><p>SIMD（单指令多数据），或者我们称之为向量处理单元。 每次都会处理一小批数据，而不是单个数据。 这需要我们统一变换循环体中的数据访问模式，以便 LLVM 后端可以将其lower到SIMD。</p></li>
</ol>
<p>实际上，本教程中使用的所有方法都是这个 <a class="reference external" href="https://github.com/flame/how-to-optimize-gemm">repo</a> 中提到的技巧的一个子集。 其中一些已经被TVM自动抽象所采用，但由于 TVM 的限制，其中一些不能简单地应用。</p>
<div class="section" id="preparation-and-performance-baseline">
<h3>准备和基线（Baseline）<a class="headerlink" href="#preparation-and-performance-baseline" title="永久链接至标题">¶</a></h3>
<p>我们首先收集有关矩阵乘法的 <cite>numpy</cite> 实现的性能数据。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">tvm</span>
<span class="kn">import</span> <span class="nn">tvm.testing</span>
<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">te</span>
<span class="kn">import</span> <span class="nn">numpy</span>

<span class="c1"># The size of the matrix</span>
<span class="c1"># (M, K) x (K, N)</span>
<span class="c1"># You are free to try out different shapes, sometimes TVM optimization outperforms numpy with MKL.</span>
<span class="n">M</span> <span class="o">=</span> <span class="mi">1024</span>
<span class="n">K</span> <span class="o">=</span> <span class="mi">1024</span>
<span class="n">N</span> <span class="o">=</span> <span class="mi">1024</span>

<span class="c1"># The default tensor data type in tvm</span>
<span class="n">dtype</span> <span class="o">=</span> <span class="s2">&quot;float32&quot;</span>

<span class="c1"># You will want to adjust the target to match any CPU vector extensions you</span>
<span class="c1"># might have. For example, if you&#39;re using using Intel AVX2 (Advanced Vector</span>
<span class="c1"># Extensions) ISA for SIMD, you can get the best performance by changing the</span>
<span class="c1"># following line to ``llvm -mcpu=core-avx2``, or specific type of CPU you use.</span>
<span class="c1"># Recall that you&#39;re using llvm, you can get this information from the command</span>
<span class="c1"># ``llc --version`` to get the CPU type, and you can check ``/proc/cpuinfo``</span>
<span class="c1"># for additional extensions that your processor might support.</span>

<span class="n">target</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span><span class="p">(</span><span class="n">target</span><span class="o">=</span><span class="s2">&quot;llvm&quot;</span><span class="p">,</span> <span class="n">host</span><span class="o">=</span><span class="s2">&quot;llvm&quot;</span><span class="p">)</span>
<span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">device</span><span class="p">(</span><span class="n">target</span><span class="o">.</span><span class="n">kind</span><span class="o">.</span><span class="n">name</span><span class="p">,</span> <span class="mi">0</span><span class="p">)</span>

<span class="c1"># Random generated tensor for testing</span>
<span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">numpy</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">rand</span><span class="p">(</span><span class="n">M</span><span class="p">,</span> <span class="n">K</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">numpy</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">rand</span><span class="p">(</span><span class="n">K</span><span class="p">,</span> <span class="n">N</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>

<span class="c1"># Repeatedly perform a matrix multiplication to get a performance baseline</span>
<span class="c1"># for the default numpy implementation</span>
<span class="n">np_repeat</span> <span class="o">=</span> <span class="mi">100</span>
<span class="n">np_running_time</span> <span class="o">=</span> <span class="n">timeit</span><span class="o">.</span><span class="n">timeit</span><span class="p">(</span>
    <span class="n">setup</span><span class="o">=</span><span class="s2">&quot;import numpy</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;M = &quot;</span> <span class="o">+</span> <span class="nb">str</span><span class="p">(</span><span class="n">M</span><span class="p">)</span> <span class="o">+</span> <span class="s2">&quot;</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;K = &quot;</span> <span class="o">+</span> <span class="nb">str</span><span class="p">(</span><span class="n">K</span><span class="p">)</span> <span class="o">+</span> <span class="s2">&quot;</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;N = &quot;</span> <span class="o">+</span> <span class="nb">str</span><span class="p">(</span><span class="n">N</span><span class="p">)</span> <span class="o">+</span> <span class="s2">&quot;</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s1">&#39;dtype = &quot;float32&quot;</span><span class="se">\n</span><span class="s1">&#39;</span>
    <span class="s2">&quot;a = numpy.random.rand(M, K).astype(dtype)</span><span class="se">\n</span><span class="s2">&quot;</span>
    <span class="s2">&quot;b = numpy.random.rand(K, N).astype(dtype)</span><span class="se">\n</span><span class="s2">&quot;</span><span class="p">,</span>
    <span class="n">stmt</span><span class="o">=</span><span class="s2">&quot;answer = numpy.dot(a, b)&quot;</span><span class="p">,</span>
    <span class="n">number</span><span class="o">=</span><span class="n">np_repeat</span><span class="p">,</span>
<span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Numpy running time: </span><span class="si">%f</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="n">np_running_time</span> <span class="o">/</span> <span class="n">np_repeat</span><span class="p">))</span>

<span class="n">answer</span> <span class="o">=</span> <span class="n">numpy</span><span class="o">.</span><span class="n">dot</span><span class="p">(</span><span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">())</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Numpy running time: 0.018646
</pre></div>
</div>
<p>现在我们使用 TVM TE 编写一个基本的矩阵乘法，并验证它可以产生和 numpy 实现相同的结果。 我们还编写了一个函数来帮助我们测量调度优化的性能。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># TVM Matrix Multiplication using TE</span>
<span class="n">k</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">reduce_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">K</span><span class="p">),</span> <span class="s2">&quot;k&quot;</span><span class="p">)</span>
<span class="n">A</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">K</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;A&quot;</span><span class="p">)</span>
<span class="n">B</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">K</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;B&quot;</span><span class="p">)</span>
<span class="n">C</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="k">lambda</span> <span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">:</span> <span class="n">te</span><span class="o">.</span><span class="n">sum</span><span class="p">(</span><span class="n">A</span><span class="p">[</span><span class="n">x</span><span class="p">,</span> <span class="n">k</span><span class="p">]</span> <span class="o">*</span> <span class="n">B</span><span class="p">[</span><span class="n">k</span><span class="p">,</span> <span class="n">y</span><span class="p">],</span> <span class="n">axis</span><span class="o">=</span><span class="n">k</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;C&quot;</span><span class="p">)</span>

<span class="c1"># Default schedule</span>
<span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>
<span class="n">func</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">)</span>

<span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">numpy</span><span class="o">.</span><span class="n">zeros</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="o">=</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">func</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">answer</span><span class="p">,</span> <span class="n">rtol</span><span class="o">=</span><span class="mf">1e-5</span><span class="p">)</span>


<span class="k">def</span> <span class="nf">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="nb">vars</span><span class="p">,</span> <span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="p">,</span> <span class="n">optimization</span><span class="p">,</span> <span class="n">log</span><span class="p">):</span>
    <span class="n">func</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">)</span>
    <span class="k">assert</span> <span class="n">func</span>

    <span class="n">c</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">numpy</span><span class="o">.</span><span class="n">zeros</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="o">=</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
    <span class="n">func</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span>
    <span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">answer</span><span class="p">,</span> <span class="n">rtol</span><span class="o">=</span><span class="mf">1e-5</span><span class="p">)</span>

    <span class="n">evaluator</span> <span class="o">=</span> <span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span><span class="p">(</span><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">number</span><span class="o">=</span><span class="mi">10</span><span class="p">)</span>
    <span class="n">mean_time</span> <span class="o">=</span> <span class="n">evaluator</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">)</span><span class="o">.</span><span class="n">mean</span>
    <span class="nb">print</span><span class="p">(</span><span class="s2">&quot;</span><span class="si">%s</span><span class="s2">: </span><span class="si">%f</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="n">optimization</span><span class="p">,</span> <span class="n">mean_time</span><span class="p">))</span>
    <span class="n">log</span><span class="o">.</span><span class="n">append</span><span class="p">((</span><span class="n">optimization</span><span class="p">,</span> <span class="n">mean_time</span><span class="p">))</span>


<span class="n">log</span> <span class="o">=</span> <span class="p">[]</span>

<span class="n">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;none&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>none: 3.003668
</pre></div>
</div>
<p>让我们使用 TVM lowe函数看一下算子的中间表示（IR）和默认的调度。 请注意，该实现本质上是一个原始的矩阵乘法实现，在 A 和 B 矩阵的索引上使用三个嵌套的循环。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (x: int32, 0, 1024) {
    for (y: int32, 0, 1024) {
      C_2[((x*1024) + y)] = 0f32
      for (k: int32, 0, 1024) {
        C_2[((x*1024) + y)] = ((float32*)C_2[((x*1024) + y)] + ((float32*)A_2[((x*1024) + k)]*(float32*)B_2[((k*1024) + y)]))
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-1-blocking">
<h3>优化1：分块<a class="headerlink" href="#optimization-1-blocking" title="永久链接至标题">¶</a></h3>
<p>提高缓存命中率的一个重要技巧是分块——数据块将逐块进行计算。 块内部的内存访问是一个具有高内存局部性的小邻域。 在本教程中，我选择了 32 作为分块因子。 所以该块将填充 32 * 32 * sizeof(float) 即 4KB 到缓存中，缓存的总大小为 32KB（L1 数据缓存）</p>
<p>我们首先为`C`` 运算符创建一个默认的调度，然后基于指定的块因子对其应用 <code class="docutils literal notranslate"><span class="pre">tile</span></code> 调度源语，该源语从外层到最内层循环的结果顺序，如向量 <code class="docutils literal notranslate"><span class="pre">[x_outer,</span> <span class="pre">y_outer,</span> <span class="pre">x_inner,</span> <span class="pre">y_inner]</span></code> 所示。然后我们获得运算符输出的约分轴，并对其使用因子为4的拆分操作。这个因子不会直接影响我们现在正在进行的分块优化，但稍后我们应用矢量化会用到。</p>
<p>既然操作已经分好块了，我们可以对计算进行重排，将约分操作放到计算的最外层循环中，以保证块里面的数据都能保留在缓存中。这就完成了调度，然后我们就可以构建程序并测试与朴素调度相比的性能了。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">bn</span> <span class="o">=</span> <span class="mi">32</span>

<span class="c1"># Blocking by loop tiling</span>
<span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">yi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">tile</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">1</span><span class="p">],</span> <span class="n">bn</span><span class="p">,</span> <span class="n">bn</span><span class="p">)</span>
<span class="p">(</span><span class="n">k</span><span class="p">,)</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
<span class="n">ko</span><span class="p">,</span> <span class="n">ki</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>

<span class="c1"># Hoist reduction domain outside the blocking loop</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">ko</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">yi</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;blocking&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>blocking: 0.217251
</pre></div>
</div>
<p>通过重排计算以利用缓存，你应该会看到计算性能的显著提高。 现在，打印中间表示并将其与原始表示进行比较：</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (x.outer: int32, 0, 32) {
    for (y.outer: int32, 0, 32) {
      for (x.inner.init: int32, 0, 32) {
        for (y.inner.init: int32, 0, 32) {
          C_2[((((x.outer*32768) + (x.inner.init*1024)) + (y.outer*32)) + y.inner.init)] = 0f32
        }
      }
      for (k.outer: int32, 0, 256) {
        for (k.inner: int32, 0, 4) {
          for (x.inner: int32, 0, 32) {
            for (y.inner: int32, 0, 32) {
              C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = ((float32*)C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] + ((float32*)A_2[((((x.outer*32768) + (x.inner*1024)) + (k.outer*4)) + k.inner)]*(float32*)B_2[((((k.outer*4096) + (k.inner*1024)) + (y.outer*32)) + y.inner)]))
            }
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-2-vectorization">
<h3>优化2：向量化<a class="headerlink" href="#optimization-2-vectorization" title="永久链接至标题">¶</a></h3>
<p>另一个重要的技巧是矢量化。 当内存访问模式一致时，编译器可以检测到这种模式并将连续内存传递给向量处理器。 在 TVM 中，我们可以使用 <cite>vectorize</cite> 接口来提示编译器这种模式，利用硬件的特性。</p>
<p>在本教程中，我们选择矢量化内循环行数据，因为它已经从我们之前的优化中对缓存是友好的。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># Apply the vectorization optimization</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">yi</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;vectorization&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>

<span class="c1"># The generalized IR after vectorization</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>vectorization: 0.251549
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (x.outer: int32, 0, 32) {
    for (y.outer: int32, 0, 32) {
      for (x.inner.init: int32, 0, 32) {
        C_2[ramp((((x.outer*32768) + (x.inner.init*1024)) + (y.outer*32)), 1, 32)] = broadcast(0f32, 32)
      }
      for (k.outer: int32, 0, 256) {
        for (k.inner: int32, 0, 4) {
          for (x.inner: int32, 0, 32) {
            C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] = ((float32x32*)C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.inner*1024)) + (k.outer*4)) + k.inner)], 32)*(float32x32*)B_2[ramp((((k.outer*4096) + (k.inner*1024)) + (y.outer*32)), 1, 32)]))
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-3-loop-permutation">
<h3>优化3：循环重排<a class="headerlink" href="#optimization-3-loop-permutation" title="永久链接至标题">¶</a></h3>
<p>如果我们查看上面的 IR，我们可以看到 B 和 C 的内循环行数据都进行了向量化。接下来我们将查看 A 的访问模式。在当前调度中，A 是逐列访问的，这对缓存不友好。 如果我们改变  <cite>ki</cite> 和内轴  <cite>xi</cite> 的嵌套循环顺序，A 矩阵的访问模式会对缓存更加友好。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>
<span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">yi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">tile</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">1</span><span class="p">],</span> <span class="n">bn</span><span class="p">,</span> <span class="n">bn</span><span class="p">)</span>
<span class="p">(</span><span class="n">k</span><span class="p">,)</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
<span class="n">ko</span><span class="p">,</span> <span class="n">ki</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>

<span class="c1"># re-ordering</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">ko</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">yi</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">yi</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span>
    <span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;loop permutation&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span>
<span class="p">)</span>

<span class="c1"># Again, print the new generalized IR</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>loop permutation: 0.098122
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  for (x.outer: int32, 0, 32) {
    for (y.outer: int32, 0, 32) {
      for (x.inner.init: int32, 0, 32) {
        C_2[ramp((((x.outer*32768) + (x.inner.init*1024)) + (y.outer*32)), 1, 32)] = broadcast(0f32, 32)
      }
      for (k.outer: int32, 0, 256) {
        for (x.inner: int32, 0, 32) {
          for (k.inner: int32, 0, 4) {
            C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] = ((float32x32*)C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.inner*1024)) + (k.outer*4)) + k.inner)], 32)*(float32x32*)B_2[ramp((((k.outer*4096) + (k.inner*1024)) + (y.outer*32)), 1, 32)]))
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-4-array-packing">
<h3>优化4：数组压缩<a class="headerlink" href="#optimization-4-array-packing" title="永久链接至标题">¶</a></h3>
<p>另一个重要的技巧是数组压缩。这个技巧是对多维数组的存储顺序进行重排，以便在将其展平并存储在一维内存中后按顺序访问。</p>
<img alt="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/array-packing.png" class="align-center" src="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/array-packing.png" />
<p>如上图所示，计算分块之后，我们可以观察到展平后B的数组访问模式不是连续的。我们期望通过一些变换后可以获得一个连续的访问模式。通过将一个 <code class="docutils literal notranslate"><span class="pre">[16][16]</span></code> 数组重排为 <code class="docutils literal notranslate"><span class="pre">[16/4][16][4]</span></code> 数组，当我们从压缩的数组中获得对应的值时，B的访问模式是连续的。</p>
<p>为了实现这一点，我们必须从一个新的默认schedule开始，考虑到B的新包装。值得花点时间对此进行评论：TE是一种功能强大的用于编写优化算子的表达性语言，但它通常需要一些底层算法、数据结构、以及您正在编写的硬件target。在本教程的后面，我们将讨论让TVM承担一些选项。不管怎样，让我们继续优化新的schedule。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># We have to re-write the algorithm slightly.</span>
<span class="n">packedB</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">((</span><span class="n">N</span> <span class="o">/</span> <span class="n">bn</span><span class="p">,</span> <span class="n">K</span><span class="p">,</span> <span class="n">bn</span><span class="p">),</span> <span class="k">lambda</span> <span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">z</span><span class="p">:</span> <span class="n">B</span><span class="p">[</span><span class="n">y</span><span class="p">,</span> <span class="n">x</span> <span class="o">*</span> <span class="n">bn</span> <span class="o">+</span> <span class="n">z</span><span class="p">],</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;packedB&quot;</span><span class="p">)</span>
<span class="n">C</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span>
    <span class="p">(</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span>
    <span class="k">lambda</span> <span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">:</span> <span class="n">te</span><span class="o">.</span><span class="n">sum</span><span class="p">(</span><span class="n">A</span><span class="p">[</span><span class="n">x</span><span class="p">,</span> <span class="n">k</span><span class="p">]</span> <span class="o">*</span> <span class="n">packedB</span><span class="p">[</span><span class="n">y</span> <span class="o">//</span> <span class="n">bn</span><span class="p">,</span> <span class="n">k</span><span class="p">,</span> <span class="n">tvm</span><span class="o">.</span><span class="n">tir</span><span class="o">.</span><span class="n">indexmod</span><span class="p">(</span><span class="n">y</span><span class="p">,</span> <span class="n">bn</span><span class="p">)],</span> <span class="n">axis</span><span class="o">=</span><span class="n">k</span><span class="p">),</span>
    <span class="n">name</span><span class="o">=</span><span class="s2">&quot;C&quot;</span><span class="p">,</span>
<span class="p">)</span>

<span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>

<span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">yi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">tile</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">1</span><span class="p">],</span> <span class="n">bn</span><span class="p">,</span> <span class="n">bn</span><span class="p">)</span>
<span class="p">(</span><span class="n">k</span><span class="p">,)</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
<span class="n">ko</span><span class="p">,</span> <span class="n">ki</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>

<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">ko</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">yi</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">yi</span><span class="p">)</span>

<span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">z</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">z</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">x</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;array packing&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>

<span class="c1"># Here is the generated IR after array packing.</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>array packing: 0.113916
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32x32), float32x32, [32768]), storage_scope = global {
    for (x: int32, 0, 32) &quot;parallel&quot; {
      for (y: int32, 0, 1024) {
        packedB[ramp(((x*32768) + (y*32)), 1, 32)] = (float32x32*)B_2[ramp(((y*1024) + (x*32)), 1, 32)]
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.inner.init: int32, 0, 32) {
          C_2[ramp((((x.outer*32768) + (x.inner.init*1024)) + (y.outer*32)), 1, 32)] = broadcast(0f32, 32)
        }
        for (k.outer: int32, 0, 256) {
          for (x.inner: int32, 0, 32) {
            for (k.inner: int32, 0, 4) {
              C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] = ((float32x32*)C_2[ramp((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.inner*1024)) + (k.outer*4)) + k.inner)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + (k.inner*32)), 1, 32)]))
            }
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-5-optimizing-block-writing-through-caching">
<h3>优化 5：通过缓存优化块写入<a class="headerlink" href="#optimization-5-optimizing-block-writing-through-caching" title="永久链接至标题">¶</a></h3>
<p>至此，我们所有的优化都专注于高效访问和计算来自”A”和”B”矩阵的数据，以计算”C”矩阵。阻止优化后，算子会将结果逐块写入”C”块，并且访问模式不是连续的。我们可以通过使用顺序缓存阵列，使用”cache_write”，”compute_at”和”展开”的组合来保持块结果，并在所有块结果准备好时写入”C”。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>

<span class="c1"># Allocate write cache</span>
<span class="n">CC</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_write</span><span class="p">(</span><span class="n">C</span><span class="p">,</span> <span class="s2">&quot;global&quot;</span><span class="p">)</span>

<span class="n">xo</span><span class="p">,</span> <span class="n">yo</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">yi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">tile</span><span class="p">(</span><span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">C</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span><span class="p">[</span><span class="mi">1</span><span class="p">],</span> <span class="n">bn</span><span class="p">,</span> <span class="n">bn</span><span class="p">)</span>

<span class="c1"># Write cache is computed at yo</span>
<span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">],</span> <span class="n">yo</span><span class="p">)</span>

<span class="c1"># New inner axes</span>
<span class="n">xc</span><span class="p">,</span> <span class="n">yc</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>

<span class="p">(</span><span class="n">k</span><span class="p">,)</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
<span class="n">ko</span><span class="p">,</span> <span class="n">ki</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">ko</span><span class="p">,</span> <span class="n">xc</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">yc</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">unroll</span><span class="p">(</span><span class="n">ki</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">CC</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">yc</span><span class="p">)</span>

<span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">z</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">z</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">x</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;block caching&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span><span class="p">)</span>

<span class="c1"># Here is the generated IR after write cache blocking.</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>block caching: 0.117208
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32x32), float32x32, [32768]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global {
    for (x: int32, 0, 32) &quot;parallel&quot; {
      for (y: int32, 0, 1024) {
        packedB[ramp(((x*32768) + (y*32)), 1, 32)] = (float32x32*)B_2[ramp(((y*1024) + (x*32)), 1, 32)]
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.c.init: int32, 0, 32) {
          C.global[ramp((x.c.init*32), 1, 32)] = broadcast(0f32, 32)
        }
        for (k.outer: int32, 0, 256) {
          for (x.c: int32, 0, 32) {
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[(((x.outer*32768) + (x.c*1024)) + (k.outer*4))], 32)*(float32x32*)packedB[ramp(((y.outer*32768) + (k.outer*128)), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 1)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 32), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 2)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 64), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 3)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 96), 1, 32)]))
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="optimization-6-parallelization">
<h3>优化 6：并行化<a class="headerlink" href="#optimization-6-parallelization" title="永久链接至标题">¶</a></h3>
<p>到目前为止，我们的计算只设计为使用单个核心。几乎所有的现代处理器都有多个内核，计算可以同时运行计算。最终优化是利用线程级并行。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># parallel</span>
<span class="n">s</span><span class="p">[</span><span class="n">C</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">xo</span><span class="p">)</span>

<span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">z</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">z</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">packedB</span><span class="p">]</span><span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">x</span><span class="p">)</span>

<span class="n">evaluate_operation</span><span class="p">(</span>
    <span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;mmult&quot;</span><span class="p">,</span> <span class="n">optimization</span><span class="o">=</span><span class="s2">&quot;parallelization&quot;</span><span class="p">,</span> <span class="n">log</span><span class="o">=</span><span class="n">log</span>
<span class="p">)</span>

<span class="c1"># Here is the generated IR after parallelization.</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">B</span><span class="p">,</span> <span class="n">C</span><span class="p">],</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>parallelization: 0.048986
primfn(A_1: handle, B_1: handle, C_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32x32), float32x32, [32768]), storage_scope = global {
    for (x: int32, 0, 32) &quot;parallel&quot; {
      for (y: int32, 0, 1024) {
        packedB[ramp(((x*32768) + (y*32)), 1, 32)] = (float32x32*)B_2[ramp(((y*1024) + (x*32)), 1, 32)]
      }
    }
    for (x.outer: int32, 0, 32) &quot;parallel&quot; {
      allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global;
      for (y.outer: int32, 0, 32) {
        for (x.c.init: int32, 0, 32) {
          C.global[ramp((x.c.init*32), 1, 32)] = broadcast(0f32, 32)
        }
        for (k.outer: int32, 0, 256) {
          for (x.c: int32, 0, 32) {
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[(((x.outer*32768) + (x.c*1024)) + (k.outer*4))], 32)*(float32x32*)packedB[ramp(((y.outer*32768) + (k.outer*128)), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 1)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 32), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 2)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 64), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 3)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 96), 1, 32)]))
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="summary-of-matrix-multiplication-example">
<h3>矩阵乘法示例总结<a class="headerlink" href="#summary-of-matrix-multiplication-example" title="永久链接至标题">¶</a></h3>
<p>在仅使用 18 行代码应用上述简单优化后，我们生成的代码可以开始使用数学内核库 （MKL） 接近”numpy”的性能。由于我们一直在记录我们工作时的表现，我们可以比较结果。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">baseline</span> <span class="o">=</span> <span class="n">log</span><span class="p">[</span><span class="mi">0</span><span class="p">][</span><span class="mi">1</span><span class="p">]</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="s2">&quot;Operator&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="s2">&quot;Timing&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="s2">&quot;Performance&quot;</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">)))</span>
<span class="k">for</span> <span class="n">result</span> <span class="ow">in</span> <span class="n">log</span><span class="p">:</span>
    <span class="nb">print</span><span class="p">(</span>
        <span class="s2">&quot;</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="se">\t</span><span class="si">%s</span><span class="s2">&quot;</span>
        <span class="o">%</span> <span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="nb">str</span><span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">1</span><span class="p">])</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">),</span> <span class="nb">str</span><span class="p">(</span><span class="n">result</span><span class="p">[</span><span class="mi">1</span><span class="p">]</span> <span class="o">/</span> <span class="n">baseline</span><span class="p">)</span><span class="o">.</span><span class="n">rjust</span><span class="p">(</span><span class="mi">20</span><span class="p">))</span>
    <span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>        Operator                  Timing             Performance
            none            3.0036678463                     1.0
        blocking     0.21725099509999998     0.07232856834273993
   vectorization            0.2515494919     0.08374743972102824
loop permutation     0.09812208580000001     0.03266742223873704
   array packing     0.11391647170000001     0.03792578857889544
   block caching     0.11720774190000001     0.03902153896422992
 parallelization            0.0489857072    0.016308629884073878
</pre></div>
</div>
<p>请注意，网页上的输出反映了非排他性 Docker 容器上的运行时间，应视为不可靠。我们非常鼓励自己运行教程，观察 TVM 的性能收益，并仔细研究每个示例，了解矩阵乘法操作的迭代改进。</p>
</div>
</div>
<div class="section" id="final-notes-and-summary">
<h2>最终说明总结<a class="headerlink" href="#final-notes-and-summary" title="永久链接至标题">¶</a></h2>
<p>如前所述，如何使用 TE 和调度原始信息进行优化可能需要了解基础架构和算法。但是，TE 旨在作为更复杂的算法的基础，这些算法可以搜索潜在的优化。有了您从此介绍到 TE 的知识，我们现在可以开始探索 TVM 如何自动实现计划优化过程。</p>
<p>本教程使用矢量添加和矩阵乘法示例提供了 TVM 张量表达 （TE） 工作流程的演练。一般工作流程是</p>
<ul class="simple">
<li><p>通过一系列算子描述您的计算。</p></li>
<li><p>描述我们如何使用调度原语去计算。</p></li>
<li><p>编译到我们想要的目标函数。</p></li>
<li><p>可选，保存稍后要加载的功能。</p></li>
</ul>
<p>即将推出的教程扩展到矩阵乘法示例，并显示如何使用可调参数构建矩阵乘法和其他操作的通用模板，从而允许您自动优化特定平台的计算。</p>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-tutorial-tensor-expr-get-started-py">
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../_downloads/40a01cffb015a67aaec0fad7e27cf80d/tensor_expr_get_started.py"><code class="xref download docutils literal notranslate"><span class="pre">下载</span> <span class="pre">Python</span> <span class="pre">源码:</span> <span class="pre">tensor_expr_get_started.py</span></code></a></p>
</div>
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../_downloads/4459ebf5b03d332f7f380abdaef81c05/tensor_expr_get_started.ipynb"><code class="xref download docutils literal notranslate"><span class="pre">下载</span> <span class="pre">Jupyter</span> <span class="pre">notebook:</span> <span class="pre">tensor_expr_get_started.ipynb</span></code></a></p>
</div>
</div>
<p class="sphx-glr-signature"><a class="reference external" href="https://sphinx-gallery.github.io">Gallery generated by Sphinx-Gallery</a></p>
</div>
</div>


           </div>
           
          </div>
          

<footer>

    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="autotvm_matmul_x86.html" class="btn btn-neutral float-right" title="Optimizing Operators with Schedule Templates and AutoTVM" accesskey="n" rel="next">下一个 <span class="fa fa-arrow-circle-right"></span></a>
      
      
        <a href="autotvm_relay_x86.html" class="btn btn-neutral float-left" title="Compiling and Optimizing a Model with the Python Interface (AutoTVM)" accesskey="p" rel="prev"><span class="fa fa-arrow-circle-left"></span> 上一个</a>
      
    </div>

<div id="button" class="backtop"><img src="../_static//img/right.svg" alt="backtop"/> </div>
<section class="footerSec">
    <div class="footerHeader">
      <ul class="d-flex align-md-items-center justify-content-between flex-column flex-md-row">
        <li class="copywrite d-flex align-items-center">
          <h5 id="copy-right-info">© 2020 Apache Software Foundation | All right reserved</h5>
        </li>
      </ul>

    </div>

    <ul>
      <li class="footernote">Copyright © 2020 The Apache Software Foundation. Apache TVM, Apache, the Apache feather, and the Apache TVM project logo are either trademarks or registered trademarks of the Apache Software Foundation.</li>
    </ul>

</section>
</footer>
        </div>
      </div>

    </section>

  </div>
  

    <script src="https://cdnjs.cloudflare.com/ajax/libs/popper.js/1.12.9/umd/popper.min.js" integrity="sha384-ApNbgh9B+Y1QKtv3Rn7W3mgPxhU9K/ScQsAP7hUibX39j7fakFPskvXusvfa0b4Q" crossorigin="anonymous"></script>
    <script src="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/js/bootstrap.min.js" integrity="sha384-JZR6Spejh4U02d8jOt6vLEHfe/JQGiRRSQQxSfFWpi1MquVdAyjUar5+76PVCmYl" crossorigin="anonymous"></script>

  </body>
  <script type="text/javascript">
      jQuery(function () {
          SphinxRtdTheme.Navigation.enable(true);
      });
  </script>

  
  
    
    <!-- Theme Analytics -->
    <script>
    (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){
      (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o),
      m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m)
    })(window,document,'script','https://www.google-analytics.com/analytics.js','ga');

    ga('create', 'UA-75982049-2', 'auto');
    ga('send', 'pageview');
    </script>

    
   

</body>
</html>